#include "Kernels.h"

HostHits hostHits(1);

/*--------------------------------------------------------------------------------------
 *    Function:  wordCompare
 * Description:  Compara duas palavras em CUDA
 *--------------------------------------------------------------------------------------
 */
__host__ __device__
int wordCompare(Word w1, Word w2)
{
    for(uint i = 0; i < _WD_SIZE; i++)
        if(w1[i] != w2[i] || w1[i] == '\0' || w2[i] == '\0')
            return w1[i] - w2[i];
    return 0;
}

/*--------------------------------------------------------------------------------------
 *    Function:  cudaSetupKernel
 * Description:  Inicializa os estados do CURAND
 *--------------------------------------------------------------------------------------
 */
__global__
void cudaSetupKernel(curandState* state, int seed)
{
    int id = threadIdx.x + blockIdx.x * blockDim.x;

    // Toda thread recebem sementes diferentes
    curand_init(seed + id, 0, 0, &state[id]);
}

/*--------------------------------------------------------------------------------------
 *    Function:  cudaGenerateWord
 * Description:  Gera uma palavra aleatória em CUDA
 *--------------------------------------------------------------------------------------
 */
__device__
void cudaGenerateWord(Word genWord, uint* accProb, curandState* state)
{
	// ID global da thread
    int id = threadIdx.x + blockIdx.x * blockDim.x;

    // Estado local
    curandState localState = state[id];
    // Gera um número aleatório
    uint r = curand(&localState);

    // Escolhe um tamanho aleatório
    uint len;
    for(len = _WD_SIZE; r < accProb[len - 1]; len--);

    // Escolhe aleatoriamente os caracteres
    for(uint i = 0; i < len; i++)
        genWord[i] = curand(&localState) % 26 + 'a';

    // Insere '\0'
    genWord[len] = '\0';

    state[id] = localState;
}

/*--------------------------------------------------------------------------------------
 *    Function:  cudaCheckWord
 * Description:  Gera uma palavra aleatória em CUDA
 *--------------------------------------------------------------------------------------
 */
__device__
void cudaCheckWord(Word genWord, WordPair* dict, uint dictSize, uint* hits)
{
    // Inicializa a busca binária
    uint a = 0, b = dictSize - 1, m;

    // Enquanto for um intervalo válido
    while(a <= b)
    {
        // Pega o meio
        m = (a + b) / 2;

        // Compara com a palavra do meio
        int cmp = wordCompare(genWord, dict[m].word);

        // Se for menor
        if(cmp < 0)
            b = m - 1;
        // Se for maior
        else if(cmp > 0)
            a = m + 1;
        // Se acertou
        else
        {
            if(!atomicCAS(&(dict[m].hasFound), 0, 1))
                atomicInc(hits, dictSize + 1);
            break;
        }
    }
}

/*--------------------------------------------------------------------------------------
 *    Function:  cudaDoWork
 * Description:  Função do kernel
 *--------------------------------------------------------------------------------------
 */
__global__
void cudaCountHits(WordPair *dict, uint dictSize, uint* hits, uint* accProb, curandState* state)
{
    // Gera uma palavra aleatória
    Word genWord;
    cudaGenerateWord(genWord, accProb, state);

    // Verifica quantas acertaram
    cudaCheckWord(genWord, dict, dictSize, hits);
}

/*--------------------------------------------------------------------------------------
 *    Function:  callSetupKernel
 * Description:  Chama a inicilização dos estados CURAND e copia hosts para o device
 *--------------------------------------------------------------------------------------
 */
void callSetupKernel(uint blocks, uint threads, DevState& state)
{
	cudaSetupKernel<<<blocks, threads>>>(state.data().get(), time(NULL));
}

/*--------------------------------------------------------------------------------------
 *    Function:  callGenerateWord
 * Description:  Chama a geração de palavras CUDA
 *--------------------------------------------------------------------------------------
 */
int callCountHits(uint blocks, uint threads, DevWordPair &devDict, DevHits &devHits,
        DevProb &devProb, DevState &devState)
{
    for(uint i = 0; i < BLOCK_SIZE; i += blocks * threads)
    {
        cudaCountHits<<<blocks, threads>>>(devDict.data().get(), devDict.size(), devHits.data().get(),
            devProb.data().get(), devState.data().get());
    }

	thrust::copy(devHits.begin(), devHits.end(), hostHits.begin());
	return hostHits[0];
}
